
#ifndef _CG_GRID_H
#define _CG_GRID_H
#include "maca_cooperative_groups.h"

#include "info.h"
#include "__clang_maca_atomic_functions.h"

_CG_BEGIN_NAMESPACE
namespace details
{
    _CG_STATIC_QUALIFIER void bar_flush(volatile unsigned int *addr)
    {
        /*Todo: need compiler to create this api*/
        __threadfence();
    }

    _CG_STATIC_QUALIFIER unsigned int atomic_add(volatile unsigned int *addr, unsigned int val)
    {
        unsigned int old;
        old = atomicAdd((unsigned int *)addr, val);
        return old;
    }

    _CG_STATIC_QUALIFIER unsigned int get_group_mask(unsigned int thread_rank, unsigned int num_warps)
    {
        /*Todo: need compielr to create this api.*/
        return 0;
    }
    _CG_STATIC_QUALIFIER void sync_warps(volatile barrier_t *arrived, unsigned int thread_rank, unsigned int num_warps)
    {
        /*Todo: need compiler to create this api.*/
    }

    _CG_STATIC_QUALIFIER bool sync_warps_last_releases(volatile barrier_t *arrived, unsigned int thread_rank, unsigned int num_warps)
    {
        /*Todo: need compiler to create this api.*/
        return false;
    }

    _CG_STATIC_QUALIFIER void sync_warps_release(volatile barrier_t *arrived, bool is_master, unsigned int thread_rank, unsigned int num_warps)
    {
        /*Todo: need compiler to create this api.*/
    }

    _CG_STATIC_QUALIFIER void sync_warps_arrive(volatile barrier_t *arrived, unsigned int thread_rank, unsigned int num_warps)
    {
        /*Todo: need compiler to create this api.*/
    }

    _CG_STATIC_QUALIFIER void sync_warps_arrive_release(volatile barrier_t *arrived, unsigned int thread_rank, unsigned int num_warps)
    {
        /*Todo: need compiler to create this api.*/
    }

    _CG_STATIC_QUALIFIER void sync_warps_wait(volatile barrier_t *arrived, unsigned int thread_rank)
    {
        /*Todo: need compiler to create this api.*/
    }

    _CG_STATIC_QUALIFIER void sync_warps_wait_for_release(
        volatile barrier_t *arrived,
        bool is_master,
        unsigned int thread_rank,
        unsigned int num_warps)
    {
        /*Todo: need compiler to create this api.*/
    }

    enum wait_for_warps_kind
    {
        wait_for_all_other_warps,
        wait_for_specific_warp
    };

    template <wait_for_warps_kind Kind>
    _CG_QUALIFIER void sync_warps_wait_for_warps(
        unsigned int wait_warp_id, volatile barrier_t *arrived, unsigned int thread_rank, unsigned int num_warps);

    template <>
    _CG_QUALIFIER void sync_warps_wait_for_warps<wait_for_all_other_warps>(
        unsigned int wait_warp_id,
        volatile barrier_t *arrived,
        unsigned int thread_rank,
        unsigned int num_warps)
    {
        /*Todo: need compiler to create this api.*/
    }

    template <>
    _CG_QUALIFIER void sync_warps_wait_for_warps<wait_for_specific_warp>(
        unsigned int wait_warp_id,
        volatile barrier_t *arrived,
        unsigned int thread_rank,
        unsigned int num_warps)
    {
        /*Todo: need compiler to create this api.*/
    }
}
_CG_END_NAMESPACE
#endif //_CG_GRID_H